// textures;
%s

struct Params
{
%s
};


class Shader
{
public:
// params
float time;
%s

// init params
%s

// random generator (using simple LCG for now, TODO: use a better one)
unsigned int seed;
GPU float rand() { seed = 1103515245u*seed + 12345u; return (seed + 1) * 2.328306435454494e-10f; }

// the shader itself
#line 0
%s


#ifdef EVAL_RGBA
// volume accumulation using evalRgba
GPU float3 evalRayRgba(float3 S, float3 E, float T)
{
    float rayLength = length(E - S);
    float3 rayDirection = normalize(E - S);
    float3 result = make_float3(0);
    float seeThrough = 1;
    for (float t = 0 /*rand() * 0.001f*/; t < rayLength; t += 0.001f)
    {
        float3 pos = S + t * rayDirection;
        float4 rgba = clamp(evalRgba(pos, rayDirection, T), 0, 1);
        if (rgba.w == 0) continue;
        result += seeThrough * rgba.w * make_float3(rgba);
        seeThrough *= 1 - rgba.w;
        if (seeThrough < 0.01f) break;
    }
    return result;
}
#endif
};


GPU float2 intersectSlab(float p, float d, float2 slab)
{
    if (fabs(d) < 0.0001f) return make_float2(-INF, INF);

    float x1 = (slab.x - p) / d;
    float x2 = (slab.y - p) / d;

    if (x1 <= x2) return make_float2(x1, x2);
        else return make_float2(x2, x1);
}


GPU float2 intersectIntervals(float2 a, float2 b)
{
    if (a.x > b.x)
    {
        float2 temp = a; a = b; b = temp;
    }
    
    if (b.x > a.y) return make_float2(INF, -INF);
    return make_float2(b.x, min(a.y, b.y));
}


GPU float2 intersectUnitCube(float3 p, float3 d)
{
    float2 slab = make_float2(-1, 1);

    float2 tx = intersectSlab(p.x, d.x, slab);
    float2 ty = intersectSlab(p.y, d.y, slab);
    float2 tz = intersectSlab(p.z, d.z, slab);
    
    return intersectIntervals(tx, intersectIntervals(ty, tz));
}


class Rgba
{
public:
    unsigned char r, g, b, a;
    
    GPU Rgba(float3 rgb, float scale, float invGamma)
    {
        rgb = clamp(rgb * scale, 0.0f, 1.0f);
        r = (unsigned char)(pow(rgb.x, invGamma) * 255);
        g = (unsigned char)(pow(rgb.y, invGamma) * 255);
        b = (unsigned char)(pow(rgb.z, invGamma) * 255);
        a = 255;
    }
    
    GPU Rgba()
    {
        r = g = b = 0;
        a = 1;
    }

};


extern "C" {

__global__ void run(
    float time,
    int imgWidth,
    int imgHeight,
    Frame* frame,
    float zoom,
    float dist,
    Params* params,
    float3* hdrOut,
    Rgba* ldrOut)
{
    // get image pixel coordinates
    int j = threadIdx.x + blockDim.x * blockIdx.x;
    int i = threadIdx.y + blockDim.y * blockIdx.y;
    if (i >= imgHeight || j >= imgWidth) return;
    
    // get normalized coordinates (vertical image axis maps to [-1,1])
    float x = j * 2 - imgWidth;
    float y = i * 2 - imgHeight;
    x /= imgHeight; y /= imgHeight;
    
    // setup ray
    float3 rayOrigin, rayDirection;
    if (dist <= 0)
    {
        rayOrigin = frame->getPointToWorld(make_float3(x / zoom, y / zoom, 0));
        rayDirection = frame->getVectorToWorld(make_float3(0, 0, 1));
    }
    else
    {
        rayOrigin = frame->origin;
        rayDirection = normalize(frame->getVectorToWorld(make_float3(x / zoom, y / zoom, dist)));
    }
    
    // clip the ray to the [-1,1]^3 cube
    float2 interval = intersectUnitCube(rayOrigin, rayDirection);
    float3 color;
    
    Shader shader;
    shader.init(params);
    
    if (interval.x == INF) color = make_float3(0);
    else
    {
        float3 start = rayOrigin + interval.x * rayDirection;
        float3 end = rayOrigin + interval.y * rayDirection;
        
        // initialize random generator
        shader.seed = tex2D(randomTexture, (j & 255) + 0.5f, (i & 255) + 0.5f);
        
        // evaluate ray
#ifdef EVAL_RGBA
        color = shader.evalRayRgba(start, end, time);
#else        
        color = shader.evalRay(start, end, time);
#endif
    }
    
    if (hdrOut != 0)
        hdrOut[i*imgWidth + j] = color;
    if (ldrOut != 0)
        ldrOut[i*imgWidth + j] = Rgba(color, 1, 1.0f / 2.2f);
}

}

